//
//  ChannelsSendingTester2.cpp
//  GPUProcessingTesting
//
//  Created by kadlubek47 on 17.06.2013.
//  Copyright (c) 2013 kadlubek47. All rights reserved.
//

#include "ChannelsSendingTester2.h"



#pragma mark Create, Destroy

const char ChannelsSendingTester2::kernelPath_[] = "/Users/kadlubeczek47/programowanie/Praca Inzynierska/parallel-computing-sound-signals_svn_tmp/trunk/GPUProcessingTesting/GPUProcessingTesting/simply_kernels.cl";

ChannelsSendingTester2::ChannelsSendingTester2(): OpenCLProcessorBase((char*)kernelPath_) {
    createKernels();
}

ChannelsSendingTester2::~ChannelsSendingTester2() {
    if(!createdBuffers_) {
        clReleaseMemObject(bufferIn_);
        clReleaseMemObject(bufferOut_);
        clReleaseMemObject(buffer_channelsInfo_);
    }
    
    clReleaseKernel(inOutKernel_);
}

void ChannelsSendingTester2::createKernels() {
    inOutKernel_ = clCreateKernel(program_, "in_out2", &lastCommandStatus_);
}



#pragma mark manyChannel main_function

void ChannelsSendingTester2::manyChannels(float* channels, unsigned int channelsNum, unsigned int channelLength, float* destiny, double* timeOfCallCommandsFun, double* timeOfFlushAndFinish) {
    Timer timer;
    
    //Calling commands functions
    timer.start();
    
    if( (channelsNum != channelsNum_) || (channelLength != channelLength_) ) {
        channelsNum_ = channelsNum;
        channelLength_ = channelLength;
        createBuffers();
    }
    
    loadBlocksOfChannelsIntoGPU(channels);
    executeKernel();
    getBlocksOfChannelsFromGPU(destiny);
    
    timer.stop();
    *timeOfCallCommandsFun += timer.getTime();
    //.
    
    //Flushing and Finishing
    timer.start();
    
    clFlush(cmdQueue_);
    clFinish(cmdQueue_);
    
    timer.stop();
    *timeOfFlushAndFinish += timer.getTime();
    //.
}



#pragma mark manyChannel subordinate_functions

void ChannelsSendingTester2::createBuffers() {
    
    //Delete previous buffers with different size
    if( createdBuffers_ ) {
        clReleaseMemObject(bufferIn_);
        clReleaseMemObject(bufferOut_);
    }
    //.
    
    //Create channels info buffer
    if( ! createdBuffers_ ) {
        buffer_channelsInfo_ = clCreateBuffer(
                                       context_,
                                       CL_MEM_READ_ONLY,
                                       2 * sizeof(cl_uint),
                                       NULL,
                                       &lastCommandStatus_);
    }
    
    //Create in out buffers
    bufferIn_ = clCreateBuffer(
                                       context_,
                                       CL_MEM_READ_ONLY,
                                       channelLength_ * channelsNum_ * sizeof(float),
                                       NULL,
                                       &lastCommandStatus_);
    bufferOut_ = clCreateBuffer(
                                       context_,
                                       CL_MEM_WRITE_ONLY,
                                       channelLength_ * channelsNum_ * sizeof(float),
                                       NULL,
                                       &lastCommandStatus_);
    //.
    
    //Write data to channels info buffer
    cl_uint* info = new cl_uint[2];
    info[0] = (cl_uint)channelsNum_;     //< info[1] == channelsNum_
    info[1] = (cl_uint)channelLength_;   //< info[0] == channelLength_
    lastCommandStatus_ = clEnqueueWriteBuffer(
                                      cmdQueue_,
                                      buffer_channelsInfo_,
                                      CL_FALSE,
                                      0,
                                      sizeof(cl_uint) * 2,
                                      info,
                                      0,
                                      NULL,
                                      NULL);
    //.
    
    //Set kernel arguments
    lastCommandStatus_  = clSetKernelArg(
                                    inOutKernel_,
                                    0,
                                    sizeof(cl_mem),
                                    &bufferIn_);
    
    lastCommandStatus_ |= clSetKernelArg(
                                    inOutKernel_,
                                    1,
                                    sizeof(cl_mem),
                                    &bufferOut_);
    
    lastCommandStatus_ |= clSetKernelArg(
                                    inOutKernel_,
                                    2,
                                    sizeof(cl_mem),
                                    &buffer_channelsInfo_ );
    //.
    
    //Set flag of creation
    createdBuffers_ = true;
    //.
}

void ChannelsSendingTester2::loadBlocksOfChannelsIntoGPU(float* channels) {
    
    lastCommandStatus_ = clEnqueueWriteBuffer(
                                    cmdQueue_,
                                    bufferIn_,
                                    CL_FALSE,
                                    0,
                                    channelLength_ * channelsNum_ * sizeof(float),
                                    channels,
                                    0,
                                    NULL,
                                    NULL);
    
}

void ChannelsSendingTester2::executeKernel() {
    size_t globalWorkSize[1];
    globalWorkSize[0] = channelLength_ ;
    
    lastCommandStatus_ = clEnqueueNDRangeKernel(
                                    cmdQueue_,
                                    inOutKernel_,
                                    1,
                                    NULL,
                                    globalWorkSize,
                                    NULL,
                                    0,
                                    NULL,
                                    NULL);
    
}

void ChannelsSendingTester2::getBlocksOfChannelsFromGPU (float* destiny) {
    lastCommandStatus_ = clEnqueueReadBuffer(
                                    cmdQueue_,
                                    bufferOut_,
                                    CL_FALSE,
                                    0,
                                    channelLength_ * channelsNum_ * sizeof(float),
                                    destiny,
                                    0,
                                    NULL,
                                    NULL);
}